#include<stdio.h>
#include"fftz.h"

static const int N = 16777216;
static const int NPart = 4096;
static const int NNow = 64;
FftzEnv cl;
cl_mem buffer0 = NULL;
cl_mem buffer1 = NULL;
cl_mem buffer2 = NULL;
cl_kernel kernel = NULL;

int getKernel(FftzEnv* cl, char* source_str, char* fun_name, cl_kernel* kernel);
int showKernelInfo(FftzEnv* cl, cl_kernel kernel);
void release();

int main(int argc, char** argv){
	size_t gws[1];
	size_t lws[1];
	cl_event event;
	cl_ulong tqueue, tstart, tend;
	int time, mint;
	int i;
	int err;

	if(initFftzEnv(&cl) < 0){
		printf("Init FftzEvn failed, exitting..\n");
		return -1;
	}

	buffer0 = clCreateBuffer(cl.context, CL_MEM_READ_WRITE, sizeof(float) * N * 2, NULL, &err);
	if(err < 0){
		printf("Creating buffer0 failed, exitting..\n");
		release();
		return err;
	}
	buffer1 = clCreateBuffer(cl.context, CL_MEM_READ_WRITE, sizeof(float) * N * 2, NULL, &err);
	if(err < 0){
		printf("Creating buffer1 failed, exitting..\n");
		release();
		return err;
	}
	buffer2 = clCreateBuffer(cl.context, CL_MEM_READ_WRITE, sizeof(float) * NPart * 4, NULL, &err);
	if(err < 0){
		printf("Creating buffer2 failed, exitting..\n");
		release();
		return err;
	}
	
	err = getKernel(&cl, strKernelR4, "gpuFftFour", &kernel);
	if(err < 0){
		printf("Getting kernel failed, exitting..\n");
		release();
		return err;
	}
	showKernelInfo(&cl, kernel);

	err += clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer0);
	err += clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer1);
	err += clSetKernelArg(kernel, 2, sizeof(cl_mem), &buffer2);
	err += clSetKernelArg(kernel, 3, sizeof(int), &N);
	err += clSetKernelArg(kernel, 4, sizeof(int), &NNow);
	err += clSetKernelArg(kernel, 5, sizeof(int), &NPart);
	if(err < 0){
		printf("Failed to set kernel args, exitting..\n");
		release();
		return err;
	}

	gws[0] = N / 4;
	lws[0] = 256;
	printf("running time:\n  ");
	mint = 1000000;
	for(i = 0; i < 10; i++){
		err = clEnqueueNDRangeKernel(cl.queue, kernel, 1, NULL, gws, lws, 0, NULL, &event);
		if(err < 0){
			printf("Failed to enqueue kernel, exitting..\n");
			release();
			return err;
		}
		clFinish(cl.queue);
		clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &tstart, NULL);
		clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &tend, NULL);
		time = (tend - tstart) / 1000;	//us
		clReleaseEvent(event);
		printf("%8d", time);
		if(time < mint)	mint = time;
	}
	printf("\n  min time: %d\n", mint);

	release();
	return 0;
}

int getKernel(FftzEnv* cl, char* source_str, char* fun_name, cl_kernel* kernel){
	cl_program myprogram;
	size_t program_build_log_size;
	char* program_build_log;
	int err;

	//compile program
	myprogram = clCreateProgramWithSource(cl->context, 1, (const char**)&source_str, NULL, &err);
	err += clBuildProgram(myprogram, 1, cl->device, NULL, NULL, NULL);
	if(err < 0){
		//print program compiling compiling information
		clGetProgramBuildInfo(myprogram, cl->device[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &program_build_log_size);
		program_build_log = (char*)malloc(sizeof(char) * program_build_log_size);
		clGetProgramBuildInfo(myprogram, cl->device[0], CL_PROGRAM_BUILD_LOG, program_build_log_size, program_build_log, NULL);
		printf("%s\n", program_build_log);
		free(program_build_log);

		printf("Error: Compiling program failed!\n");
		return err;
	}

	//get kernel
	*kernel = NULL;
	*kernel = clCreateKernel(myprogram, fun_name, &err);
	if(err < 0){
		printf("Failed to create kernel %s, function exiting...\n", fun_name);
		clReleaseProgram(myprogram);
		return err;
	}

	clReleaseProgram(myprogram);
	return 0;
}

int showKernelInfo(FftzEnv* cl, cl_kernel kernel){
	size_t groupSize;
	int err;
	
	err = clGetKernelWorkGroupInfo(kernel, cl->device[0], CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(groupSize), &groupSize, NULL);
	if(err < 0){
		printf("Failed to get CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE information, function exiting..\n");
		return -1;
	}
	printf("  preferred work group size multiple: %ld\n", groupSize);
	
	err = clGetKernelWorkGroupInfo(kernel, cl->device[0], CL_KERNEL_WORK_GROUP_SIZE, sizeof(groupSize), &groupSize, NULL);
	if(err < 0){
		printf("Failed to get CL_KERNEL_WORK_GROUP_SIZE information, function exiting..\n");
		return -1;
	}
	printf("  preferred work group size: %ld\n", groupSize);
	
	return 0;
}

void release(){
	if(buffer0 != NULL)	clReleaseMemObject(buffer0);
	if(buffer1 != NULL) clReleaseMemObject(buffer1);
	if(buffer2 != NULL) clReleaseMemObject(buffer2);
	if(kernel != NULL) clReleaseKernel(kernel);
	releaseFftzEnv(&cl);
}
